#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <sys/fcntl.h>
#include <sys/mman.h>
#include <unistd.h>
#include <string.h>
#include <cuda.h>
#include "scanImpl.cu"

#define BI 1000000000
#define HSIZE (1024*1024*10)

#  define CUDA_SAFE_CALL_NO_SYNC( call) do {                                \
    cudaError err = call;                                                    \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)
#  define CUDA_SAFE_CALL( call) do {                                        \
    CUDA_SAFE_CALL_NO_SYNC(call);                                            \
    cudaError err = cudaThreadSynchronize();                                 \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)

__device__ __host__ static int gpu_atoi(const char *string, int size)
{
	int i, length = 0;
	i=0;
	while(length <size){
		i=10*i + (*string - '0');
		string++;
		length ++;
	}

	return i;
}

__device__ __host__ static unsigned int StringHash(const char* s, int size)
{
    unsigned int hash = 0;
    int c, i=0;

    while(i<size)
    {
	c = *s++;
        hash = ((hash << 5) + hash) ^ c;
	i++;
    }

    return hash;
}

__global__ void count_hash_num(char * dim, int tupleNum,int *num, int keySize, int *key){
	int stride = blockDim.x * gridDim.x;
	int offset = blockIdx.x * blockDim.x + threadIdx.x;
	//int rnum = (keySize + sizeof(int)-1)/sizeof(int);
	//int end = (tupleNum * keySize + sizeof(int)-1)/sizeof(int);
	//char buf[128];

	for(int i=offset;i<tupleNum;i+=stride){
		//int start = i/blockDim.x*blockDim.x*keySize/sizeof(int);
		//int start = (i * keySize)/sizeof(int);
		//for(int j=0;j<rnum && (start+j)<end;j++){
	//		((char4 *)buf)[j] = ((char4*)dim)[start +j]; 
	//	}
		//int pos = i* keySize %sizeof(int);
		int hKey = StringHash(dim + i*keySize,keySize) % HSIZE;
		key[i] = hKey;
		num[hKey] = 1;
	}
}

__global__ void count_group_num(int *num, int tupleNum, int *totalCount){
	int stride = blockDim.x * gridDim.x;
	int offset = blockIdx.x * blockDim.x + threadIdx.x;
	int localCount = 0;

	for(int i=offset; i<tupleNum; i+= stride){
		if(num[i] == 1){
			localCount ++;
		}
	}

	atomicAdd(totalCount,localCount);
}

__global__ void agg_cal(char *dim, char *col, int *key, int tupleNum, int *psum, char * bucket, int keySize, int colNum, int*keyBool){

	int stride = blockDim.x * gridDim.x;
	int index = blockIdx.x * blockDim.x + threadIdx.x;

	for(int i=index;i<tupleNum;i+=stride){

                int hKey = key[i];
                int offset = psum[hKey] * (colNum*sizeof(int)+(keySize+sizeof(int)-1)/sizeof(int)*sizeof(int));
		if(keyBool[hKey]==0){
			memcpy(bucket+offset, dim + i*keySize, keySize);
			keyBool[hKey]=1;
		}
		offset += (keySize +sizeof(int)-1)/sizeof(int) * sizeof(int);
		for(int j=0;j<colNum;j++){
			atomicAdd(&((int *)(bucket+offset))[j],((int *)col)[j*tupleNum +i]);
		}
	}
}

int main(int argc, char * argv[]){
	int *gbCol, gbKey;

	int * gpu_hashNum;
	char *gpu_bucket, * gpu_key;

	struct timespec start,end;
	float gpuTime;
	char **gbTable;
	char *keyTable;

	if(argc!=5){
		printf("error: gbTablePrefix colNum gbkeyWidth distinct\n");
		exit(-1);
	}

	int * tmp;
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&tmp, sizeof(int)));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(tmp));

	int colNum = atoi(argv[2]);
	int keySize = atoi(argv[3]);
	int distinct = atoi(argv[4]);

	gbCol = (int *) malloc(sizeof(int) * colNum);
	gbTable = (char **)malloc(sizeof(char *) * colNum);

	if(!gbCol || !gbTable){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	char buf[32] = {0};
	for(int i=0;i<colNum;i++){
		memset(buf,0,sizeof(buf));
		sprintf(buf,"%s%d",argv[1],i);
		gbCol[i] = open(buf,O_RDONLY);
	}

	int colSize = lseek(gbCol[0],0,SEEK_END);
	int tupleNum = colSize/sizeof(int);

	for(int i=0;i<colNum;i++){
		gbTable[i] = (char *) malloc(colSize);
		char *gbT = (char *)mmap(0,colSize,PROT_READ,MAP_SHARED,gbCol[i],0);
		memcpy(gbTable[i],gbT,colSize);
		munmap(gbT,colSize);
	}

	memset(buf,0,sizeof(buf));
	sprintf(buf,"gbKey");
	gbKey = open(buf, O_RDONLY);
	
	keyTable = (char *) malloc(tupleNum * keySize); 
	char *keyT = (char *)mmap(0,tupleNum * keySize,PROT_READ,MAP_SHARED,gbKey,0);
	memcpy(keyTable,keyT,tupleNum * keySize);
	munmap(keyT,tupleNum * keySize);

	char * host_result = (char *)malloc(2*distinct * (colNum *sizeof(int) + keySize));
	memset(host_result, 0, 2*distinct*(colNum * sizeof(int) + keySize) );

	dim3 grid(512);
	dim3 block(32);
	char * gpu_col;

/*
 * 	build hash table on GPU
 */

	clock_gettime(CLOCK_REALTIME,&start);

	cudaEvent_t startGPU, stopGPU;
	cudaEventCreate(&startGPU);
	cudaEventCreate(&stopGPU);
	cudaEventRecord(startGPU,0);
	int *gpu_psum;

	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpu_col, colSize*colNum));

	int offset = 0;

	for(int i=0;i<colNum;i++){
		CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(gpu_col+offset, gbTable[i],colSize,cudaMemcpyHostToDevice));
		offset += colSize;
	}

	int totalSize = (tupleNum * keySize + sizeof(int)-1 )/sizeof(int)*sizeof(int);
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpu_key,totalSize));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(gpu_key,keyTable,tupleNum * keySize,cudaMemcpyHostToDevice));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpu_hashNum,sizeof(int)*HSIZE));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemset(gpu_hashNum,0,sizeof(int)*HSIZE));

	int * gpuHashKey;
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpuHashKey, sizeof(int)*tupleNum));

	count_hash_num<<<grid,block>>>(gpu_key,tupleNum,gpu_hashNum,keySize, gpuHashKey);

	int * gpuGroupNum, groupNum;

	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpuGroupNum, sizeof(int)));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemset(gpuGroupNum, 0, sizeof(int)));

	count_group_num<<<grid,block>>>(gpu_hashNum, HSIZE, gpuGroupNum);

	CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(&groupNum, gpuGroupNum, sizeof(int),cudaMemcpyDeviceToHost));

	printf("%d\n",groupNum);

	CUDA_SAFE_CALL(cudaMalloc((void**)&gpu_psum,HSIZE*sizeof(int)));

	scanImpl(gpu_hashNum,HSIZE,gpu_psum);

	int resultSize = groupNum * (colNum * sizeof(int) + (keySize+sizeof(int)-1)/sizeof(int) * sizeof(int));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpu_bucket, resultSize));

	int * keyBool;
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&keyBool,sizeof(int)*HSIZE));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemset(keyBool, 0, sizeof(int)*HSIZE));
	agg_cal<<<grid,block>>>(gpu_key,gpu_col,gpuHashKey,tupleNum,gpu_psum,gpu_bucket, keySize,colNum, keyBool);

	CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(host_result,gpu_bucket, resultSize, cudaMemcpyDeviceToHost));

	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpu_key));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpu_hashNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpu_psum));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpu_bucket));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuHashKey));

	for(int i=0;i<colNum;i++){
		close(gbCol[i]);
		free(gbTable[i]);
	}
	close(gbKey);

	cudaEventRecord(stopGPU,0);
	cudaEventSynchronize(stopGPU);
	cudaEventElapsedTime(&gpuTime,startGPU,stopGPU);
	printf("GPU total time:%lf\n",gpuTime);

	clock_gettime(CLOCK_REALTIME,&end);
	double timeE = (end.tv_sec -  start.tv_sec)* BI + end.tv_nsec - start.tv_nsec;
	printf("CPU time: %lf\n", timeE);
	return 0;
}
